Skip to content

Fp16 constant weights - #3015

Open
mloubout wants to merge 8 commits into
mainfrom
fp16-constant-weights
Open

Fp16 constant weights#3015
mloubout wants to merge 8 commits into
mainfrom
fp16-constant-weights

Conversation

@mloubout

@mloubout mloubout commented Sep 2, 2026

Copy link
Copy Markdown
Contributor

On top of #3012

TODO:

Add a knob for fp16 weights

@mloubout mloubout added the API api (symbolics, types, ...) label Sep 2, 2026
@codecov

codecov Bot commented Sep 2, 2026

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 76.92308% with 3 lines in your changes missing coverage. Please review.
✅ Project coverage is 79.00%. Comparing base (b330c14) to head (38dbac3).

Files with missing lines Patch % Lines
devito/operator/operator.py 0.00% 2 Missing and 1 partial ⚠️
Additional details and impacted files
@@            Coverage Diff             @@
##             main    #3015      +/-   ##
==========================================
- Coverage   83.72%   79.00%   -4.73%     
==========================================
  Files         257      257              
  Lines       54822    54830       +8     
  Branches     4693     4695       +2     
==========================================
- Hits        45901    43317    -2584     
- Misses       8110    10627    +2517     
- Partials      811      886      +75     
Flag Coverage Δ
pytest-gpu-aomp-amdgpuX ?
pytest-gpu-gcc- 55.05% <30.76%> (-23.29%) ⬇️
pytest-gpu-icx- ?
pytest-gpu-nvc-nvidiaX ?

Flags with carried forward coverage won't be shown. Click here to find out more.

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

The Weights of a non-expanded derivative were always built at the default
precision, so a `float16` stencil got `float` coefficients. Every
wavefield*weight product then bound to the mixed-precision operators and
was promoted, defeating the point of the half-precision wavefield.
`_gen_value` printed the initializer with the printer's default dtype
rather than the Array's, which stamped a `float` suffix onto the entries
of a `double` Array and silently rounded them to single precision.

Route it through a new `initvalue` printer hook, which also gives the
targets a place to specialize an initializer whose type cannot be built
from a plain literal.
`_prec` floors an untyped real literal at `float32` so that an integer
default doesn't degrade the arithmetic around it. That floor also caught
`float16`, which is never a fallback but an explicit request, so every
literal in a half-precision Operator printed one type too wide.

Only apply the floor when the default is not already a real type.
The stability check sums the whole field and asks whether the result is
finite. The accumulator took the field's own dtype, so in half precision
it overflowed within a few thousand points and reported an instability
that wasn't there -- making `errctl=max`, the very option one reaches for
to diagnose a suspected instability, unusable exactly where it is needed.

Give it at least single precision.
A real literal in an otherwise integer expression is emitted at the
Operator's precision, floored at `float32` so that an integer default
does not degrade it. An Operator working in half wants that floor most
of the time -- half is a storage format, and the accuracy of the
literals is worth more than the width of the multiply -- but not always.

Give the printer a flag for it, off by default, and have `_printer`
pick up a Target's second printer where one is offered.
The same coefficients at two precisions are two different arrays, but
neither `__eq__` nor `_hashable_content` looked at the dtype, so the
first one built answered for both. An Operator asking for its weights in
one precision would be handed whichever an earlier Operator had cached.

Compare and hash on it. The name goes in rather than the type itself,
which does not order and so cannot be sorted alongside the rest.
Whether an Operator working in half also computes in half decides what
is calculated, not how quickly: the literals and the FD coefficients are
rounded to three decimal digits. That is a mathematical choice, so it
belongs with `interp-mode` in `sym_opt` rather than among the codegen
options, and is validated and defaulted alongside it.
@mloubout
mloubout force-pushed the fp16-constant-weights branch from 75d6b31 to 9759188 Compare September 3, 2026 12:10

@FabioLuporini FabioLuporini left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

IIRC (a may be wrong and confusing it with the other PRO PR) there may have been leftover tiny comments in the old OSS PR that we might be able to address here

Comment thread devito/core/operator.py Outdated

INTERP_MODE = 'direct'

HALF_ARITH = False

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

you've got to move it down below -- currently it's right in between "INTERP_MODE" and its docstring


def _hashable_content(self):
return (self.name, self.dimension, str(self.weights), self.scope)
# NOTE: `dtype` belongs here. The same coefficients at two precisions

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

this NOTE can go, it's obvious

Comment thread devito/ir/iet/visitors.py Outdated
string, delegating to the printer so that languages whose types cannot
be built from plain literals (e.g. CUDA's `__half`) can specialize it.
"""
printer = get_printer(self.printer, obj.dtype)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

this seems weird to me:

get_printer(self.printer

Comment thread devito/ir/iet/visitors.py Outdated
be built from plain literals (e.g. CUDA's `__half`) can specialize it.
"""
printer = get_printer(self.printer, obj.dtype)
return printer.initvalue(init, obj.dtype)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

why, instead of this new private method, don't you just pass self.ccode(obj) and let the printer handle the whole printing of the Array? ie u let it return c.Initializer(...)

Comment thread devito/ir/cgen/printer.py Outdated
def _print_ListInitializer(self, expr):
return f"{{{', '.join(self._print(i) for i in expr.params)}}}"

def initvalue(self, init, dtype):

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

if this really is necessary (see other comments below) and if it really needs to be a public method, it should be put among the public methods, not among the private ones

@mloubout
mloubout force-pushed the fp16-constant-weights branch from fd41c9e to 9f3c81d Compare September 4, 2026 14:15
An Array initializer already reached the printer, via `ccode` on a
`ListInitializer`; what it did not do was reach it at the Array's own
precision.  The elements were printed with the Operator's settings, so a
narrow Array sitting in an Operator whose arithmetic is left at the
default width had its entries emitted at the wider type.

Pass the Array's dtype to `ccode`, as `Expression` already does, rather
than route the initializer around the printer through an `initvalue`
hook of its own.  A target that needs to spell its literals differently
overrides `_print_ListInitializer`, which is the ordinary extension
point.

With the precision now correct at the point of printing, `_prec` no
longer needs to be told whether the arithmetic was narrowed on purpose:
a real literal takes the precision it is being printed at, and the
`float32` floor applies only where that is not itself a float.  That is
the same value as before for every dtype other than `float16`.
@mloubout
mloubout force-pushed the fp16-constant-weights branch from 9f3c81d to 38dbac3 Compare September 4, 2026 14:17
@mloubout mloubout added the no-pro-trigger Skip the devitopro submodule update on merge label Sep 4, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

API api (symbolics, types, ...) no-pro-trigger Skip the devitopro submodule update on merge

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants